{# template imput - SpConf, DnConf, SpBt, SP, ET #}
{% iclude 'CoopConfig.cuh.jinja' without context %}

template <typename DenseF, typename T, typename IT, int SpVlen, int SpIters>
class SparseIterator{
  
  typedef typename llvlib::VectTypes<IT, SpVlen>::type IVT;
  typedef typename llvlib::VectTypes<T, SpVlen>::type VT;
  typedef llvlib::Vect<IT, SpVlen> IVect;
  typedef llvlib::Vect<T, SpVlen> VVect;

  static constexpr int SpTileWidth = {{ SpConf.xdim }}*SpVlen;
  static constexpr int TileRows = {{ SpConf.ydim }};
  
  template <typename message_t>
  static __device__ __forceinline__ void cache_for_dense(message_t* buffer, message_t val)
  {
    {% if SpConf.xdim != DnConf.xdim %}
    if ( !{{ xid(SpConf) }} )
    {
      buffer[{{ yid(SpConf) }}] = val;
    }
    {% endif %}
  }

  template <typename message_t>
  static __device__ __forceinline__ void retrive_for_dense(message_t* buffer, message_t& val, int idx)
  {
    {% if SpConf.xdim != DnConf.xdim %}
    val = buffer[idx];
    {% endif %}
  }

public:
  {# typedef typename buffer_t<SpBt>::template Tile<IT, SpConf::ydim, SpTileWidth> IdxBuffer; #}
  {# typedef typename buffer_t<SpBt>::template Tile<T, SpConf::ydim, SpTileWidth> ValBuffer; #}
  
  {% if spbt == 'shared' %}
  typedef typename Tile<IT, {{SpConf.ydim}}, >
  {% elif spbt == 'local' %}
  {% else %}
  {% endif %}


  static __device__ __forceinline__ bool vote(const bool pred)
  {
    bool ret = false;
    if ( DnConf::ydim >= SpConf::ydim ) { // multi-group against one row
      ret = pred;
    } else if (DnConf::xdim == 32) {
      ret = __ballot_sync(0xffffffff, pred);
    } else if (DnConf::xdim < 32) {  // one group against multi-row
      const uint32_t mask = ((uint32_t)1<<DnConf::xdim)-1;
      const uint32_t shift = threadIdx.x & (31-DnConf::xdim+1);
      ret = __ballot_sync(0xffffffff, pred) & (mask << shift);
    } else {
      __shared__ bool trans_buffer[SpConf::ydim];
      if ( !SpConf::xid() )
      {
        trans_buffer[SpConf::yid()] = pred;
      }
      __syncthreads();
      for ( int i=0; i<SpPerDn; ++i )
      {
        ret |= trans_buffer[SpPerDn*DnConf::yid()+i];
      }
    }
    return ret;
  }


  static __device__ __forceinline__ void proceed(
      IT* es_ptr, T* vs_ptr, T* src_data, 
      const int block_rows, const int block_cols, int row_nz,
      const int ld_src, const int col_base, int align_prefix,
      // llvlib::Tile<int, TileRows, SpTileWidth>& sparse_indx, llvlib::Tile<T, TileRows, SpTileWidth>& sparse_val,
      IdxBuffer& sparse_indx, ValBuffer& sparse_val,
      typename DenseF::Accumulator *local_sum, int* cache_row_nz, int* cache_align_prefix)
  {
    if (SP == SparseCompute::full || SP == SparseCompute::prefix)
    {
      if (row_nz > SpTileWidth)
      {
        sparse_indx.template gather_warp<IVT, SpConf>(es_ptr, block_rows);
        if (ET == EdgeType::w) sparse_val.template gather_warp<VT, SpConf>(vs_ptr, block_rows);
        cache_for_dense(cache_row_nz, row_nz);
        cache_for_dense(cache_align_prefix, align_prefix);
      } else {
        row_nz = -1;
        cache_for_dense(cache_row_nz, -1);
      }
    }
    else // suffix
    {
      // sparse_indx[SpConf::yid()][i] = es_ptr[i];
      sparse_indx.template gather_warp<IVT, SpConf>(es_ptr, block_rows, row_nz);
      if (ET == EdgeType::w) //sparse_val[SpConf::yid()][i] = vs_ptr[i];
        sparse_val.template gather_warp<VT, SpConf>(vs_ptr, block_rows, row_nz);
      // if (blockIdx.x == 0 && blockIdx.y == 0) printf("loading %d from %lx to buffer[%d,%d]\n", sparse_indx[SpConf::yid()][i], es_ptr+i, SpConf::yid(), i);
        
      cache_for_dense(cache_row_nz, row_nz);
      cache_for_dense(cache_align_prefix, align_prefix); 
    
      // sparse_indx.template gather_warp<int, SpConf>(es_ptr, block_rows, nz_row);
      // if (ET == EdgeType::w) sparse_val.template gather_warp<T, SpConf>(vs_ptr, block_rows, nz_row);
    }
    if ( CoopCoordinator::LHS_major || SpBt != BufferType::none ) SpConf::sync();

    if ( block_cols > 0 ) for (int r=0; r<SpPerDn; ++r)
    {
      const int local_row = SpPerDn*DnConf::yid() + r;
      if ( local_row >= block_rows ) break;
      
      retrive_for_dense(cache_row_nz, row_nz, local_row);
      if ( row_nz == -1 ) continue;
      
      retrive_for_dense(cache_align_prefix, align_prefix, local_row);
    
      const int tile_s = (SP == SparseCompute::full) ? 0 : align_prefix;
      const int tile_e = (SP == SparseCompute::suffix) ? row_nz : SpTileWidth;
      // const IT *loc_rows  = (SP == SparseCompute::suffix) ? loc_es_ptr[local_row] : sparse_indx[local_row];
      // const T  *loc_vals  = (SP == SparseCompute::suffix) ? loc_vs_ptr[local_row] : sparse_val[local_row];
      
      IT *loc_cols = sparse_indx[local_row];
      T  *loc_vals = sparse_val[local_row];

      const int subid = CoopCoordinator::sub_yid_L2R();
      for (int i=tile_s+subid; i<tile_e; i+=DnPerSp)
      {
        const size_t dense_row = static_cast<size_t>(loc_cols[i]);
        const T      spval     = (ET == EdgeType::w) ? loc_vals[i] : 0;

        // if (blockIdx.x == 0 && blockIdx.y == 0 && DnConf::xid() == 0)
        // printf("Thread %d,%d: loc_vals[%d~%d~%d]=%f\n", DnConf::yid(), DnConf::xid(), tile_s, i, tile_e, loc_vals[i]);
        // if (RowOffset+SpConf::yid() == 1007 && SpConf::xid() == 0) printf("local_group %d loading %d from buffer[%d,%d]\n", SpConf::yid(), dense_row, local_row, i);
        T* src_row = src_data + ld_src*dense_row + col_base;
        DenseF::compute_dense_tile(src_row, spval, local_sum[r], block_cols);
      }

    }
    if ( CoopCoordinator::LHS_major || SpBt != BufferType::none ) DnConf::sync();
  }

  static __device__ __forceinline__ void 
  writeback(const IT* row_indx, T* dst_data, const int ld_dst, 
            const int col_base, const int block_rows, const int block_cols,
            typename DenseF::Accumulator *local_sum)
  {
    for (int r=0; r<SpPerDn; ++r)
    {
      const int local_row = SpPerDn*DnConf::yid() + r;
      if ( local_row >= block_rows ) break;

      const int wb_row = SpConf::ydim*blockIdx.x + local_row;
      size_t sp_row    = static_cast<size_t>(row_indx[wb_row]);
      // size_t sp_row = SpConf::ydim*blockIdx.y + r*DnConf::ydim + DnConf::yid();
      T* dst_base = dst_data + ld_dst*sp_row + col_base;
      // if (SpConf::xid() == 0) printf("row=%d dst_ptr=%lx\n", sp_row, dst_base);
      DenseF::write_dense(dst_base, local_sum[r], block_cols);
    }
  }
};